#define L1_BUFFER_LEN ::std::uint32_t(128)

#pragma GCC optimize("O3")
__global__ void kernel_silu(float const* const __restrict dev_inp, float * const __restrict dev_out, ::std::uint32_t const nr_elems) {
    tops_dte_ctx_t ctx;
    // FIXME use dte_scope instead
    ctx.init();

    // Calculate the unique thread index within the grid
    ::std::uint32_t const thread_idx {blockIdx.x * blockDim.x + threadIdx.x};

    __valigned__ float l1_buf_dev_inp[L1_BUFFER_LEN];
    __valigned__ float l1_buf_dev_out[L1_BUFFER_LEN];
    tops::mdspan l1_buf_dev_out_view(tops::Private, l1_buf_dev_out, L1_BUFFER_LEN);

    ::std::uint32_t const elems_per_thread {nr_elems / 24};
    ::std::uint32_t const pos_end {thread_idx == 23 ? nr_elems : (thread_idx + 1) * elems_per_thread};
    for (::std::uint32_t i {thread_idx * elems_per_thread}; i < pos_end; i += L1_BUFFER_LEN) {
        ::tops::mdspan l1_buf_dev_inp_view(tops::Private, l1_buf_dev_inp, L1_BUFFER_LEN);
        ::tops::mdspan l3_dev_inp_view(tops::Global, const_cast<float*>(dev_inp + i), L1_BUFFER_LEN);
        ::tops::memcpy(ctx, l1_buf_dev_inp_view, l3_dev_inp_view);

        // Perform vectorized addition on L1 buffers
        for (::std::uint32_t j{}; j < L1_BUFFER_LEN; j += /* tops::vlength<vfloat>() = */ 32) {
            auto const data = tops::vload<vfloat>(l1_buf_dev_inp + j);
            auto const sigmoid = ::tops::vsigmoid<vfloat>(data);
            ::tops::vstore(data * sigmoid, l1_buf_dev_out + j);
        }

        ::tops::mdspan l3_dev_out_view(tops::Global, dev_out + i, L1_BUFFER_LEN);
        ::tops::memcpy(ctx, l3_dev_out_view, l1_buf_dev_out_view);
    }
}

void GCU_SILU(float * __restrict dev_inp, float * __restrict dev_out, const int nr_elems) {
    ::kernel_silu<<<2, 12>>>(dev_inp, dev_out, ::std::uint32_t(nr_elems));
}
